
#include "fun.h"
#include "hs_cuda_elemwise.cuh"

template<typename T>
struct OddEvenFunctor {
    __device__ void operator()(T x, T y, T & z1, T & z2) const
    {
        z1 = (x + y) / 2.0;
        z2 = (x - y) / 2.0;
    }
};

void testOddEven(const hs::raster::RasterData<float>& host_mat_a, const hs::raster::RasterData<float>& host_mat_b, hs::raster::RasterData<float>& host_mat_c, hs::raster::RasterData<float>& host_mat_d)
{
    /* 输入数据分配显存并拷贝数据 */
    hscuda::raster::RasterData<float> device_mat_a;
    hscuda::raster::RasterData<float> device_mat_b;
    hscuda::raster::host2device(host_mat_a.view(), device_mat_a);
    hscuda::raster::host2device(host_mat_b.view(), device_mat_b);

    /* 输出数据分配显存 */
    int size_x = device_mat_a.view().sizeX();
    int size_y = device_mat_a.view().sizeY();
    hscuda::raster::RasterData<float> device_mat_c(size_x, size_y);
    hscuda::raster::RasterData<float> device_mat_d(size_x, size_y);
    
    /* 执行并行 */
    int N = device_mat_a.view().numel();
    hs::cuda::elemwise::mimo::Op2in2out(N
        , device_mat_a.view().pData()
        , device_mat_b.view().pData()
        , OddEvenFunctor<float>()
        , device_mat_c.view().pData()
        , device_mat_d.view().pData());
    cudaDeviceSynchronize();
    
    /* 将结果保存到主机端 */
    device_mat_c.toHost(host_mat_c);
    device_mat_d.toHost(host_mat_d);
}
